Skip to content

Extend SUM_WITH_OVERFLOW reduction to signed integers and decimals - #22465

Merged
rapids-bot[bot] merged 16 commits into
NVIDIA:mainfrom
PointKernel:sum-with-overflow-reduction-extend
May 22, 2026
Merged

Extend SUM_WITH_OVERFLOW reduction to signed integers and decimals#22465
rapids-bot[bot] merged 16 commits into
NVIDIA:mainfrom
PointKernel:sum-with-overflow-reduction-extend

Conversation

@PointKernel

Copy link
Copy Markdown
Member

Description

Closes #22413

This PR drops the INT64-only guard on the SUM_WITH_OVERFLOW reduction and broadens it to all signed integer types (int8/int16/int32/int64) and decimal types (decimal32/64/128), matching the trait and the groupby path. Output preserves the input element type. Tests are unified into one typed fixture covering all newly supported types.

Checklist

  • I am familiar with the Contributing Guidelines.
  • New or existing tests cover these changes.
  • The documentation is up to date with these changes.

@PointKernel PointKernel added feature request New feature or request libcudf Affects libcudf (C++/CUDA) code. labels May 11, 2026
@copy-pr-bot

copy-pr-bot Bot commented May 11, 2026

Copy link
Copy Markdown

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

@PointKernel PointKernel self-assigned this May 11, 2026
@PointKernel PointKernel added the non-breaking Non-breaking change label May 11, 2026
@PointKernel
PointKernel marked this pull request as ready for review May 13, 2026 22:22
@PointKernel
PointKernel requested a review from a team as a code owner May 13, 2026 22:22
@PointKernel
PointKernel requested review from shrshi and ttnghia May 13, 2026 22:22
@coderabbitai

coderabbitai Bot commented May 13, 2026

Copy link
Copy Markdown

Review Change Stack

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

  • @coderabbitai resume to resume automatic reviews.
  • @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

  • ▶️ Resume reviews
  • 🔍 Trigger review
📝 Walkthrough

Walkthrough

Extends SUM_WITH_OVERFLOW from int64-only to signed integer and fixed-point types; implements a templated, type-dispatched reduction with an overflow-aware accumulator and struct result; updates docs and SPDX years; adds typed tests and error-case tests covering the new type matrix.

Changes

SUM_WITH_OVERFLOW Reduction Type Extension

Layer / File(s) Summary
Public docs and SPDX updates
cpp/include/cudf/reduction.hpp, cpp/include/cudf/reduction/detail/reduction_functions.hpp, cpp/src/reductions/reductions.cpp
Update SPDX years and expand documentation to state SUM_WITH_OVERFLOW supports signed integer and decimal (fixed-point) inputs and returns STRUCT{col.type, BOOL8}.
Type constraints and dispatcher
cpp/src/reductions/reductions.cpp, cpp/src/reductions/sum_with_overflow.cu
Broadened reduction template constraint from INT64-only to signed integral (non-bool) and fixed-point types; added constrained dispatcher and changed public sum_with_overflow to route via type_dispatcher; unsupported types call CUDF_FAIL.
Core templated reduction logic
cpp/src/reductions/sum_with_overflow.cu
Added {sum, overflow} accumulator/result types, overflow-combining operator, value and null-aware transformers, and make_sum_overflow_struct_scalar; implemented sum_with_overflow_impl<Source> handling empty/all-null, invalid init, thrust::transform_reduce reduction, and returning the {sum, overflow} struct scalar with correct validity.
Tests: typed coverage and error cases
cpp/tests/reductions/reduction_tests.cpp
Refactored monomorphic tests into ReduceWithOverflowTest<T> typed-suite over signed integers and decimal fixed-point types; added sum_overflow_rep mapping for fixed-point reps; expanded tests for normal sums, overflow directions, empty/all-null, slicing, multi-block, init-value cases; added ReduceWithOverflowErrorTest asserting std::invalid_argument for unsupported types; updated ReductionIsValidTest for decimal.

Estimated code review effort

🎯 4 (Complex) | ⏱️ ~60 minutes

Possibly related PRs

  • rapidsai/cudf#22404: JNI/Java binding changes related to exposing and testing the SUM_WITH_OVERFLOW aggregation.

Suggested reviewers

  • paul-aiyedun
🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%. Write docstrings for the functions missing them to satisfy the coverage threshold.
✅ Passed checks (4 passed)
Check name Status Explanation
Title check ✅ Passed The title clearly and concisely summarizes the main change: extending the SUM_WITH_OVERFLOW reduction to support signed integers and decimals beyond just int64.
Description check ✅ Passed The PR description relates directly to the changeset, explaining the motivation, implementation scope, and checklist completion for extending SUM_WITH_OVERFLOW support.
Linked Issues check ✅ Passed The code changes fully implement the requirements in #22413: dispatching to signed integers (int8/16/32/64) and decimals (decimal32/64/128), preserving output type, and removing the INT64-only guard while aligning with groupby.
Out of Scope Changes check ✅ Passed All changes are directly related to extending SUM_WITH_OVERFLOW support; copyright year updates, documentation updates, implementation changes, and test refactoring are all in scope.

✏️ Tip: You can configure your own custom pre-merge checks in the settings.

✨ Finishing Touches
🧪 Generate unit tests (beta)
  • Create PR with unit tests

Comment @coderabbitai help to get the list of available commands and usage tips.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 2

🧹 Nitpick comments (1)
cpp/tests/reductions/reduction_tests.cpp (1)

3476-3707: ⚡ Quick win

Add sliced and multi-block cases to the new typed suite.

The implementation now runs through column_device_view iterators and transform_reduce, but these tests only exercise small contiguous inputs. A sliced column and a > block_size input would give much better coverage for the new offset-handling and parallel-reduction paths, especially for decimal inputs.

As per coding guidelines, cpp/**/*test*.{cu,cpp}: Test functions must cover edge cases: empty input, null values, sliced columns, boundary sizes, multi-block sizes.

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/reductions/reduction_tests.cpp` around lines 3476 - 3707, Add tests
for sliced and multi-block inputs to the ReduceWithOverflowTest typed suite:
create a sliced column (use cudf::test::column_wrapper's slice or cudf::slice to
produce a column with non-zero offset) and run cudf::reduce(...) with
*cudf::make_sum_with_overflow_aggregation<reduce_aggregation>() to verify
correct sum and overflow extraction via
ReduceWithOverflowTest::extract_sum_overflow, and add a large column whose size
exceeds the backend block size (e.g., replicate a value to produce > block_size
elements) to exercise the parallel reduction path; also add analogous cases for
null-containing/sliced columns using make_null_col and for fixed_point types
(numeric::decimal32/64/128) so the typed suite covers offset handling and
multi-block reductions.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@cpp/src/reductions/sum_with_overflow.cu`:
- Around line 129-133: The code currently treats an invalid init as if no init
was supplied; change the init handling in sum_overflow_result<DeviceType>
initial_value so that when init.has_value() is true you check
init.value().get().is_valid(stream) and if that check fails set
initial_value.valid = false (and do not seed initial_value.sum), otherwise seed
initial_value.sum from the scalar and set initial_value.valid = true; this makes
SUM_WITH_OVERFLOW propagate an invalid init to an invalid result consistent with
the other reductions (refer to symbols initial_value, sum_overflow_result,
init.value().get().is_valid).
- Around line 41-58: overflow_sum_op's operator() is not associative and
performs undefined signed overflow; fix by computing the combined sum in a safe,
well-defined wider or unsigned domain and deriving the overflow flag
deterministically so combination order doesn't matter: in
overflow_sum_op::operator(), when combining two sum_overflow_result<DeviceType>
values, first promote lhs.sum and rhs.sum to a wider integer type (or their
unsigned counterpart) or use a checked-add routine to compute a full-precision
result without invoking signed overflow, set overflow = lhs.overflow ||
rhs.overflow || (result out of DeviceType bounds), and store the
clamped/truncated DeviceType sum only after that safe check; also ensure that if
lhs.overflow or rhs.overflow is already true you do not perform a signed
addition on DeviceType (use the wider/unsigned accumulator or skip addition and
propagate a canonical value) so the combiner remains associative and avoids UB.

---

Nitpick comments:
In `@cpp/tests/reductions/reduction_tests.cpp`:
- Around line 3476-3707: Add tests for sliced and multi-block inputs to the
ReduceWithOverflowTest typed suite: create a sliced column (use
cudf::test::column_wrapper's slice or cudf::slice to produce a column with
non-zero offset) and run cudf::reduce(...) with
*cudf::make_sum_with_overflow_aggregation<reduce_aggregation>() to verify
correct sum and overflow extraction via
ReduceWithOverflowTest::extract_sum_overflow, and add a large column whose size
exceeds the backend block size (e.g., replicate a value to produce > block_size
elements) to exercise the parallel reduction path; also add analogous cases for
null-containing/sliced columns using make_null_col and for fixed_point types
(numeric::decimal32/64/128) so the typed suite covers offset handling and
multi-block reductions.
🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

  • Push a commit to this branch (recommended)
  • Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 71daa890-cfd2-48af-ba12-342ac80a5967

📥 Commits

Reviewing files that changed from the base of the PR and between b262f18 and 81caf9e.

📒 Files selected for processing (3)
  • cpp/src/reductions/reductions.cpp
  • cpp/src/reductions/sum_with_overflow.cu
  • cpp/tests/reductions/reduction_tests.cpp

Comment thread cpp/src/reductions/sum_with_overflow.cu
Comment thread cpp/src/reductions/sum_with_overflow.cu Outdated
@PointKernel PointKernel moved this to Burndown in libcudf May 14, 2026
@PointKernel
PointKernel force-pushed the sum-with-overflow-reduction-extend branch from 81caf9e to f316ba0 Compare May 14, 2026 23:14

@shrshi shrshi left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Minor nits, but looks great overall!
Note (courtesy Claude): can we also update the doxygen documentation for sum_with_overflow in reduction_functions.hpp to indicate that new types are supported?

Comment thread cpp/src/reductions/sum_with_overflow.cu
Comment thread cpp/src/reductions/sum_with_overflow.cu

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 1

🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@cpp/include/cudf/reduction.hpp`:
- Around line 80-82: The documentation for the throw condition in
cudf::reduction (the comment describing std::invalid_argument for the
sum_with_overflow reduction) uses the wrong parameter name `output_dtype`;
update that identifier to the correct API parameter name `output_type` so the
doc matches the function signature and references `sum_with_overflow` and
`output_type` consistently.
🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

  • Push a commit to this branch (recommended)
  • Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 2b6d5159-8e25-4479-b1ca-a98941211889

📥 Commits

Reviewing files that changed from the base of the PR and between f316ba0 and 5f21626.

📒 Files selected for processing (3)
  • cpp/include/cudf/reduction.hpp
  • cpp/include/cudf/reduction/detail/reduction_functions.hpp
  • cpp/src/reductions/sum_with_overflow.cu
✅ Files skipped from review due to trivial changes (1)
  • cpp/include/cudf/reduction/detail/reduction_functions.hpp
🚧 Files skipped from review as they are similar to previous changes (1)
  • cpp/src/reductions/sum_with_overflow.cu

Comment thread cpp/include/cudf/reduction.hpp
Comment thread cpp/src/reductions/sum_with_overflow.cu

@bdice bdice left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

One comment (see above), otherwise LGTM.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🧹 Nitpick comments (1)
cpp/src/reductions/sum_with_overflow.cu (1)

103-107: ⚡ Quick win

Use cudf::get_current_device_resource_ref() for temporary allocations.

The overflow_scalar and children columns are temporaries—only used to construct the returned struct_scalar. Per coding guidelines, temporary memory should use cudf::get_current_device_resource_ref(), reserving the passed-in mr for the final returned memory.

♻️ Proposed fix
-  auto overflow_scalar = cudf::make_fixed_width_scalar<bool>(overflow_value, stream, mr);
+  auto overflow_scalar = cudf::make_fixed_width_scalar<bool>(overflow_value, stream, cudf::get_current_device_resource_ref());

   std::vector<std::unique_ptr<cudf::column>> children;
-  children.push_back(cudf::make_column_from_scalar(*sum_scalar, 1, stream, mr));
-  children.push_back(cudf::make_column_from_scalar(*overflow_scalar, 1, stream, mr));
+  children.push_back(cudf::make_column_from_scalar(*sum_scalar, 1, stream, cudf::get_current_device_resource_ref()));
+  children.push_back(cudf::make_column_from_scalar(*overflow_scalar, 1, stream, cudf::get_current_device_resource_ref()));

As per coding guidelines: "Allocate temporary memory using cudf::get_current_device_resource_ref(), not the passed-in MR parameter".

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/src/reductions/sum_with_overflow.cu` around lines 103 - 107, The
temporaries overflow_scalar and the temporary columns vector children (created
via cudf::make_fixed_width_scalar and cudf::make_column_from_scalar around
sum_scalar and overflow_scalar) should be allocated with
cudf::get_current_device_resource_ref() instead of the passed-in mr; change the
allocator argument for these temporary allocations to
cudf::get_current_device_resource_ref() while still using the original mr for
the final struct_scalar return so only final output uses the caller-provided
memory resource.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Nitpick comments:
In `@cpp/src/reductions/sum_with_overflow.cu`:
- Around line 103-107: The temporaries overflow_scalar and the temporary columns
vector children (created via cudf::make_fixed_width_scalar and
cudf::make_column_from_scalar around sum_scalar and overflow_scalar) should be
allocated with cudf::get_current_device_resource_ref() instead of the passed-in
mr; change the allocator argument for these temporary allocations to
cudf::get_current_device_resource_ref() while still using the original mr for
the final struct_scalar return so only final output uses the caller-provided
memory resource.

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 44779d2f-6a6b-4b4d-bc97-d4a86ddc6c3a

📥 Commits

Reviewing files that changed from the base of the PR and between a62ddd5 and 50900f2.

📒 Files selected for processing (1)
  • cpp/src/reductions/sum_with_overflow.cu

@github-actions github-actions Bot added the Java Affects Java cuDF API. label May 19, 2026
@PointKernel
PointKernel requested a review from mythrocks May 19, 2026 00:12
@PointKernel

Copy link
Copy Markdown
Member Author

@mythrocks requesting your review as you just updated the related java code.

@mythrocks

Copy link
Copy Markdown
Contributor

Apologies for the delay. (A lot of reviews at once.)
I'm perusing the changes now.

Comment thread cpp/src/reductions/sum_with_overflow.cu
@PointKernel PointKernel removed the status in libcudf May 19, 2026
Comment on lines 668 to 669
assertEquals(DType.INT64, children[0].getType());
assertEquals(DType.BOOL8, children[1].getType());

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This part of ReductionTest needs to change, to support all the new types we're adding.

At the time this was added, only INT64 reductions were possible. With added support for INT8-64, DECIMAL32-128, we can't assume that the result is INT64. :/

We will need a testSumWithOverflowSignedIntegralTypes and a testSumWithOverflowSignedDecimalTypes, for the positive cases. Those will call this helper function, to check for overflow.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Additionally, we're going to need to update the Java documentation in:

  1. Aggregation.java: Indicate that SUM_WITH_OVERFLOW accepts INT8-64, DECIMAL32-128.
  2. ReductionAggregation.java: Similar. Indicate that the output column is a STRUCT of {appropriate_result_type, overflow: BOOL}.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

cc @paul-aiyedun, for Java expertise.

I'm wondering if it makes more sense to update the Java side of this change in a separate follow-on PR.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@PointKernel: Would you mind raising a follow-up PR for the Java changes I suggested here? I don't mind that we postpone the changes. But it would be good for us to keep track.

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sure, I will do it.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I should have said follow-up issue, not PR.

It would be good for this not to drift too long, given that there are other projects now dependent on cudf-java besides spark-rapids-jni.

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks, opened #22642 to track the follow up work.

@Test
void testSumWithOverflowPositiveOverflow() {
// Long.MAX_VALUE + 1 wraps via two's complement to Long.MIN_VALUE.
// Sum is zeroed when overflow is detected; the flag is the source of truth.

@mythrocks mythrocks May 19, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is certainly the preferred behaviour, but I'm not sure that's what's happening in sum_with_overflow.cu. From there:

  return make_sum_overflow_struct_scalar<Source>(
    result.sum, result.wraps != 0, true, col.type(), stream, mr);

This looks like we're returning the sum as is.

In sum_with_overflow.cu, should we instead consider checking wraps and returning 0 conditionally?

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

(It's possible I've misunderstood this.)

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is nice with reduce but introduced some unnecessary cycles for the groupby case (see #22696) where we need to either store zero atomically or require a second pass to zero all overflowed output. With that, I'm plan to update the default behavior of both reduce and groupby to provide undeterministic sum output by bypassing the sum action whenever we see the overflow flag is set. Users should never rely on the output sum once the overflow flag is set anyway. @mythrocks, does that make sense to you?

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actually, yes. I think that should be good. We shouldn't really depend on the 0, if the overflow flag is set.

I brought it up since comment in the reduction-test indicates a zero sum. It sounded like a good idea, but maybe it's more trouble than it's worth.

Comment thread cpp/src/reductions/sum_with_overflow.cu Outdated
@PointKernel
PointKernel requested a review from mythrocks May 21, 2026 18:46
rapids-bot Bot pushed a commit that referenced this pull request May 21, 2026
Uncovered this cleanup opportunity while working on #22465.

This PR replaces cudf's custom overflow handling logic with the overflow utilities provided by CCCL. Note that although these utilities are under the `cuda::` namespace, they are available on both host and device.

Authors:
  - Yunsong Wang (https://github.com/PointKernel)

Approvers:
  - Vyas Ramasubramani (https://github.com/vyasr)
  - Muhammad Haseeb (https://github.com/mhaseeb123)
  - David Wendt (https://github.com/davidwendt)

URL: #22593

@mythrocks mythrocks left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm 👍 on the change, with the proviso that we file a follow-up issue to address the remaining Java concerns described here:
We need to update ReductionTest.java with the newly supported types for SUM_WITH_OVERFLOW.

It looks like you've already updated the documentation for the Aggregation to indicate the newly supported types, and the behaviour in the event of an overflow. Thank you for that.

@PointKernel

Copy link
Copy Markdown
Member Author

/merge

@rapids-bot
rapids-bot Bot merged commit 3cfe15b into NVIDIA:main May 22, 2026
217 of 219 checks passed
@PointKernel
PointKernel deleted the sum-with-overflow-reduction-extend branch May 22, 2026 17:00
rapids-bot Bot pushed a commit that referenced this pull request Jun 2, 2026
Originally proposed by @shrshi at #19403 (comment) and follows the subsequent discussions in #22465 (comment)

This PR updates both `groupby` and `reduce` `SUM_WITH_OVERFLOW` so that once a group's (or the reduction's) overflow flag is set, the sum is no longer maintained: `groupby` short-circuits subsequent atomic adds for that group, and `reduce` stops zeroing the wrapped result. The sum child is documented as unspecified on overflow; the boolean flag is the sole source of truth.

Note that this is an ABI-breaking change, since the expected overflowed sum value is no longer the wrapped result. However, this should not impact downstream users, as relying on overflowed sum results is not considered valid behavior.

Authors:
  - Yunsong Wang (https://github.com/PointKernel)

Approvers:
  - Nghia Truong (https://github.com/ttnghia)
  - Vukasin Milovanovic (https://github.com/vuule)
  - David Wendt (https://github.com/davidwendt)

URL: #22696
@GregoryKimball GregoryKimball removed this from libcudf Jun 3, 2026
@coderabbitai coderabbitai Bot mentioned this pull request Jun 9, 2026
3 tasks
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

feature request New feature or request Java Affects Java cuDF API. libcudf Affects libcudf (C++/CUDA) code. non-breaking Non-breaking change

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[FEA] Extend SUM_WITH_OVERFLOW reduction to signed integers and decimals

6 participants